优先级：P1
这里在包初始化阶段直接导入 block_score_fa4，会把 paddlefleet_ops.flash_mask.cute.* 的 SM10/Blackwell 依赖提前带入。paddlefleet_ops.__init__ 里只有 is_flash_mask_available() 为 true 时才加载 flash_mask，否则会通过 HardwareIncompatibleBlocker 阻断 paddlefleet_ops.flash_mask 导入；因此在 H20/A100 或非 Blackwell 环境中，只要代码执行 from paddlefleet.tilelang_ops.hysparse import block_sparse_mqa_attention / sliding_window_mqa_attention，即使不走 FA4 full layer，也会在导入阶段失败。

处理要求：请针对该评论修复并提交新的 commit。

建议让 hysparse 包级入口保持非 FA4 路径可导入，把 FA4-only 的 cute import 延迟到实际调用并加可用性检查，例如：

# hysparse/__init__.py: 不要在包初始化时 eager import block_score_fa4
from .differentiable_ops import block_sparse_mqa_attention
from .pipeline import select_topk_blocks
from .swa_attn import sliding_window_mqa_attention


def block_score_fa4_attn_fwd(*args, **kwargs):
    from .block_score_fa4 import block_score_fa4_attn_fwd as _impl

    return _impl(*args, **kwargs)

同时在 block_score_fa4.py 内部用 paddlefleet_ops.is_flash_mask_available() 做明确报错/跳过，或把 FlashMaskInfoPaddle、_flash_attn_fwd/_bwd 的导入移动到 guarded 的执行路径中。

优先级：P1
新提交 99e25e6 又把这个问题带回来了：hysparse/__init__.py 仍执行 from .block_score_fa4 import block_score_fa4_attn_fwd，而当前 block_score_fa4.py 在模块顶层导入 paddlefleet_ops.flash_mask.cute.flashmask_utils/interface。相比上一版，_import_fa4() / is_flash_mask_available() 懒加载也被删除，所以非 SM100 或无 flash_mask 的环境会在导入 paddlefleet.tilelang_ops.hysparse 时先失败，连只使用 sliding_window_mqa_attention 的路径也无法加载。

处理要求：请针对该评论修复并提交新的 commit。

请恢复 package 级入口的懒加载/可用性检查，例如：

# hysparse/__init__.py 不要在模块初始化时导入 FA4-only 依赖
def block_score_fa4_attn_fwd(*args, **kwargs):
    from .block_score_fa4 import block_score_fa4_attn_fwd as _impl

    return _impl(*args, **kwargs)

优先级：P1
_BlockScoreFA4Attn.forward() 的 Tensor 输入包含 q/k/v/block_logit/startend_row_indices，但 backward 这里只返回了 dq/dk/dv。Paddle 的 PyLayer 会按前向 Tensor 输入位置映射 backward 返回值，少了 block_logit 和 startend_row_indices 的 None 槽位会导致反向执行契约不匹配；另外当前 ctx.save_for_backward(q, k, v, startend_row_indices, out, lse) 在 startend_row_indices=None（该 API 默认允许）时也会把非 Tensor 存入 saved tensors。

处理要求：请针对该评论修复并提交新的 commit。

建议同时修复保存逻辑和返回槽位，形状如下：

ctx.has_mask = startend_row_indices is not None
if ctx.has_mask:
    ctx.save_for_backward(q, k, v, startend_row_indices, out, lse)
else:
    ctx.save_for_backward(q, k, v, out, lse)
ctx.needs_grad = (
    not q.stop_gradient,
    not k.stop_gradient,
    not v.stop_gradient,
)
ctx.mark_non_differentiable(lse)

...

if ctx.has_mask:
    q, k, v, startend_row_indices, out, lse = ctx.saved_tensor()
else:
    q, k, v, out, lse = ctx.saved_tensor()
    startend_row_indices = None
...
gq, gk, gv = ctx.needs_grad
return (
    dq if gq else None,
    dk if gk else None,
    dv if gv else None,
    None,  # block_logit
    None,  # startend_row_indices
)

优先级：P1
新提交 99e25e6 回退了这个修复，而且当前 forward 还新增了 block_bos 这个 Tensor 输入：ctx.save_for_backward(q, k, v, startend_row_indices, out, lse) 仍会在 startend_row_indices is None 时保存非 Tensor，backward 仍只 return dq, dk, dv。当前 Tensor 输入至少包含 q/k/v/block_logit/block_bos，有 mask 时还包含 startend_row_indices，因此反向返回槽位还缺 block_logit、block_bos（以及 mask Tensor 对应的 None）。

处理要求：请针对该评论修复并提交新的 commit。

建议恢复“只保存 Tensor + 显式 None 槽位”的形状，并把新增的 block_bos 也纳入返回契约：

ctx.save_for_backward(q, k, v, out, lse)
ctx.startend_row_indices = startend_row_indices
ctx.has_startend = startend_row_indices is not None
ctx.needs_grad = (
    not q.stop_gradient,
    not k.stop_gradient,
    not v.stop_gradient,
)

...

gq, gk, gv = ctx.needs_grad
grads = [
    dq if gq else None,
    dk if gk else None,
    dv if gv else None,
    None,  # block_logit
    None,  # block_bos
]
if ctx.has_startend:
    grads.append(None)  # startend_row_indices
return tuple(grads)

优先级：P1
这里无条件追加 block_indices，但 block_indices 只在上面的 HySparse full-attention 分支里赋值。当前分支顺序是先判断 self.recompute_core_attention and self.training，所以当训练时同时打开 enable_hy_sparse_attention=True 和 recompute_core_attention=True 时，会先走 recompute 的普通 core_attention 分支，然后执行到这里触发 UnboundLocalError。这个配置组合是现有 MLA attention 的可用训练开关，不能因为 HySparse 打开后直接崩溃。

处理要求：请针对该评论修复并提交新的 commit。

建议把 HySparse full-attention 分支放在 recompute 分支之前，或显式禁止这个组合。若选择保留 HySparse，可按下面形状调整：

if self.config.enable_hy_sparse_attention and shared_kv is not None:
    core_attn_out, block_indices = self._hy_sparse_full_attention(
        query,
        key,
        value,
        attn_mask_startend_row_indices,
    )
elif self.recompute_core_attention and self.training:
    core_attn_out = recompute(
        self.core_attention,
        ...
    )
else:
    core_attn_out = self.core_attention(...)

如果当前 FA4 full path 暂不支持该 recompute 模式，也请在进入这里前抛出明确的 ValueError，避免后续引用未定义变量。

优先级：P1
这里把原始未 padding 的 q / do 直接传入 windowed_mqa_bwd，但 kernel 网格是 T.ceildiv(seq_len, BM)，最后一个 query tile 会执行 T.copy(Q[bb, bm * BM : (bm + 1) * BM, ...]) 和 T.copy(dO[...])。当 seq_len % block_M != 0 时，例如 s=130 且 auto-fit 选到 block_M=64，最后一个 tile 会从 [128, 192) 读 Q/dO，超过真实 seq_len=130；当前只给 bos/eos 的 BlockRange 做了 padding，写回 dQ 虽有 guard，但越界读已经发生，训练反向会有 illegal memory access 或未定义结果风险。现有 test_hysparse_windowed_swa.py 的 backward case 使用 s=128，没有覆盖这个尾块。

处理要求：请针对该评论修复并提交新的 commit。

建议在 kernel 内按行 guard 加载 Q/dO（或在 host 侧同步 padding q/do/lse/delta 后再 trim dq）。直接在 kernel 里改的形状如下：

for i, d in T.Parallel(BM, D):
    row = bm * BM + i
    in_range = row < seq_len
    safe_row = T.if_then_else(in_range, row, 0)
    Q_shared[i, d] = T.if_then_else(
        in_range, Q[bb, safe_row, bh, d], T.cast(0, dtype)
    )

for i, d in T.Parallel(BM, D_v):
    row = bm * BM + i
    in_range = row < seq_len
    safe_row = T.if_then_else(in_range, row, 0)
    dO_shared[i, d] = T.if_then_else(
        in_range, dO[bb, safe_row, bh, d], T.cast(0, dtype)
    )

同时补一个 s % block_M != 0 的 sliding_window_mqa_attention backward 测试，例如 s=130 或显式 block_M=64 的路径，防止尾块回归。

优先级：P1
这里把 MoE 判断从 issubclass(..., MoELayer) 收窄成了精确相等，导致现有 Gemma4MoELayer(MoELayer) 走到 unknown MLP 分支，additional_mlp_kwargs 不再传 pg_collection。gpt_layer_specs.py 的 attention_layer_type == "gemma4" 当前就是 LayerSpec(layer=Gemma4MoELayer, ...)，而 Gemma4MoELayer.__init__ 会把 pg_collection 继续传给基类和 router；不传会破坏 Gemma4 MoE 构建/并行配置。

处理要求：请针对该评论修复并提交新的 commit。

建议恢复子类判断：

已在当前 head 复查，MoE 判断已恢复为 isinstance(..., type) + issubclass(..., MoELayer)，Gemma4MoELayer 这类 MoELayer 子类会继续拿到 pg_collection。这条问题我视为已修复。

优先级：P1
这个提交删除了 TransformerLayer._compute_act_offload_kwargs()，但仓库现有 tests/single_card_tests/ai_edited_test/transformer/test_ai_transformer_layer.py::TestDecoderlayerActOffloadSettings 仍直接调用 TransformerLayer._compute_act_offload_kwargs(fake)；当前 head 会变成 AttributeError。同时 base forward() 里也移除了 offload_kwargs = self._compute_act_offload_kwargs() 和 recompute(..., **offload_kwargs)，会让 decoderlayer_act_offload_settings 在 full/recovery recompute 路径上静默失效。

处理要求：请针对该评论修复并提交新的 commit。

这里需要同时恢复方法和 recompute 透传，形状如下：

def _compute_act_offload_kwargs(self):
    decoderlayer_act_offload_settings = self.config.get(
        "decoderlayer_act_offload_settings", {"type": "", "value": ""}
    ) or {"type": "", "value": ""}
    ...
    return offload_kwargs

...

offload_kwargs = self._compute_act_offload_kwargs()
outputs = recompute(
    self._forward_impl,
    ...,
    **offload_kwargs,
)

优先级：P1
这个 hunk 删除了 Gemma4TransformerLayerSublayersSpec 和 Gemma4TransformerLayer，但当前代码仍然引用它们：src/paddlefleet/models/gpt/gemma4_layer_specs.py 直接 from paddlefleet.transformer.transformer_layer import Gemma4TransformerLayer，gpt_layer_specs.py 在 attention_layer_type == "gemma4" 路径导入 Gemma4TransformerLayerSublayersSpec，tests/single_card_tests/transformer/test_gemma4.py 也覆盖这些符号。当前 head 下 Gemma4 layer spec 构建会在导入阶段失败。

处理要求：请针对该评论修复并提交新的 commit。

请恢复这些定义，或把所有引用和测试同步迁移到新的定义位置；需要保留的接口形状至少包括：

@dataclass
class Gemma4TransformerLayerSublayersSpec(TransformerLayerSublayersSpec):
    post_self_attn_layernorm: LayerSpec | type = IdentityOp
    pre_mlp_layernorm: LayerSpec | type = IdentityOp
    post_mlp_layernorm: LayerSpec | type = IdentityOp


class Gemma4TransformerLayer(TransformerLayer):
    def __init__(...):
        ...

    def _forward_impl(...):
        ...

已在当前 head 复查，Gemma4TransformerLayerSublayersSpec 和 Gemma4TransformerLayer 已恢复到 transformer_layer.py，gemma4_layer_specs.py / gpt_layer_specs.py 的导入路径也重新可用。这条问题我视为已修复。

优先级：P1
新提交 99e25e6 又把这个问题带回来了：hysparse/__init__.py 仍执行 from .block_score_fa4 import block_score_fa4_attn_fwd，而当前 block_score_fa4.py 在模块顶层导入 paddlefleet_ops.flash_mask.cute.flashmask_utils/interface。相比上一版，_import_fa4() / is_flash_mask_available() 懒加载也被删除，所以非 SM100 或无 flash_mask 的环境会在导入 paddlefleet.tilelang_ops.hysparse 时先失败，连只使用 sliding_window_mqa_attention 的路径也无法加载。

处理要求：请针对该评论修复并提交新的 commit。

请恢复 package 级入口的懒加载/可用性检查，例如：

# hysparse/__init__.py 不要在模块初始化时导入 FA4-only 依赖
def block_score_fa4_attn_fwd(*args, **kwargs):
    from .block_score_fa4 import block_score_fa4_attn_fwd as _impl

    return _impl(*args, **kwargs)

优先级：P1
新提交 99e25e6 回退了这个修复，而且当前 forward 还新增了 block_bos 这个 Tensor 输入：ctx.save_for_backward(q, k, v, startend_row_indices, out, lse) 仍会在 startend_row_indices is None 时保存非 Tensor，backward 仍只 return dq, dk, dv。当前 Tensor 输入至少包含 q/k/v/block_logit/block_bos，有 mask 时还包含 startend_row_indices，因此反向返回槽位还缺 block_logit、block_bos（以及 mask Tensor 对应的 None）。

处理要求：请针对该评论修复并提交新的 commit。

建议恢复“只保存 Tensor + 显式 None 槽位”的形状，并把新增的 block_bos 也纳入返回契约：

ctx.save_for_backward(q, k, v, out, lse)
ctx.startend_row_indices = startend_row_indices
ctx.has_startend = startend_row_indices is not None
ctx.needs_grad = (
    not q.stop_gradient,
    not k.stop_gradient,
    not v.stop_gradient,
)

...

gq, gk, gv = ctx.needs_grad
grads = [
    dq if gq else None,
    dk if gk else None,
    dv if gv else None,
    None,  # block_logit
    None,  # block_bos
]
if ctx.has_startend:
    grads.append(None)  # startend_row_indices
return tuple(grads)

优先级：P1

新提交删除了 MTPEmbeddingLayer 的 spec 生成和 mtp_embedding=mtp_embedding_spec 注入。当前 GPTSublayersSpec 仍有 mtp_embedding 字段，GPTModel.get_layer_desc_list() 只有在 spec.mtp_embedding 存在时才会插入 MTPEmbeddingLayer；而 MultiTokenPredictionLayer.forward() 在 enable_mtp_magic_send=True 时会要求 dict_args["mtp_input_embeds"]，否则直接抛出 RuntimeError("MTPEmbeddingLayer may not have been executed.")。因此 enable_mtp_magic_send=True && num_nextn_predict_layers > 0 的 GPT spec 会构建出没有 MTP re-embedding 层的 pipeline，MTP forward 会失败。

处理要求：请针对该评论修复并提交新的 commit。

请恢复 magic-send 的 embedding spec 和注入，形状如下：

from paddlefleet.models.gpt.mtp_embedding_layer import MTPEmbeddingLayer

...
mtp_embedding_spec = None
if config.enable_mtp_magic_send and config.num_nextn_predict_layers > 0:
    mtp_embedding_spec = LayerSpec(
        layer=MTPEmbeddingLayer,
        extra_kwargs={"config": config},
    )

...
GPTSublayersSpec(
    ...
    mtp=mtp_layers_spec,
    mtp_embedding=mtp_embedding_spec,
    mtp_lm_head=mtp_lm_head_spec,
    ...
)

已在当前 head 复查，MTPEmbeddingLayer 的 import、mtp_embedding_spec 构造以及 GPTSublayersSpec 的 mtp_embedding 注入都已恢复，这条问题我视为已修复。

优先级：P1
HySparseTransformerLayer 这里仍只把 input_ids 传给 _forward_mlp，没有沿用基类的 origin_input_ids 透传。MoELayer.forward / TopKRouter 在 gpt_model_use_experimental_version=True 时会用 origin_input_ids 计算 aux/z loss 的原始 token mask；当前 HySparse + MoE + experimental dataflow 下会退回截断后的 input_ids，padding/MTP token 的分母会错。

处理要求：请针对该评论修复并提交新的 commit。

请和基类保持一致，在 recompute 参数、_forward_impl 签名和 _forward_mlp 调用中都传递该参数，形状如下：

outputs = recompute(
    self._forward_impl,
    ...,
    input_ids=dict_args.get("input_ids", None),
    origin_input_ids=dict_args.get("origin_input_ids", None),
    shared_key=dict_args.get("shared_key", None),
    ...,
)


def _forward_impl(..., input_ids=None, origin_input_ids=None, ...):
    ...
    output = self._forward_mlp(
        hidden_states,
        input_ids=input_ids,
        origin_input_ids=origin_input_ids,
    )

🟡 建议 这里是 HySparse MQA 生产前向路径的配置校验，不能依赖 assert。python -O 会移除该检查，错误的 rope_type 或 packed_seq_params 会继续进入 RoPE/TileLang/DSA kernel，最终表现为难定位的 shape/kernel 错误。建议改成显式 if ...: raise ValueError(...)，和上面的 CP/TP ValueError 保持一致。

🟡 建议 这个测试名和断言看起来想覆盖 MQA forward/backward，但 layer_number=0 + window_attn_skip_freq=2 会让 self.is_swa=False，因此 MQASelfAttention.is_mqa=False，第 1658 行直接走 super().forward() 的 MLA fallback。也就是说当前测试没有覆盖新增 SWA MQA 主路径、sliding_window_mqa_attention、shared_kv 消费或 sparse branch。建议把真正的 MQA 回归构造成 SWA 层（如 layer_number=1）并提供 shared_kv/backend skip，或把此用例改名为 fallback 对齐并另补真实 MQA 用例。